/*
 * cudanDynamicGlobalMemory.cpp
 *
 *  Created on: Jan 8, 2010
 *      Author: daniel
 */

#ifdef __CDT_PARSER__
#define __global__
#define __device__
#define __host__
#define __shared__
#endif

namespace CUDAN
{
	// Default Constructor.
	template<typename T>
	DynamicGlobalMemory<T>::DynamicGlobalMemory() : numthreads(0), size(0)
	{
#ifdef _DEBUG
		std::cout << "DynamicGlobalMemory<T>::DynamicGlobalMemory()" << std::endl;
#endif
		// Allocate space on the device for the current indexes.
		cudaMalloc((void**)&this->currentIndex, sizeof(int));

		// Set the indexes to 0.
		this->resetIndex();
	}

	// Construct a DynamicGlobalMemory.
	// Set the number of threads and their allocated size.
	template<typename T>
	DynamicGlobalMemory<T>::DynamicGlobalMemory(int numthreads, int size) : numthreads(numthreads), size(size)
	{
#ifdef _DEBUG
		std::cout << "DynamicGlobalMemory<T>::DynamicGlobalMemory(int numthreads, int size)" << std::endl;
#endif
		// Allocate space on the device for the current indexes.
		cudaMalloc((void**)&this->currentIndex, sizeof(int) * numthreads);

		// Set the indexes to 0.
		this->resetIndex();

		// Set the size of the pre-allocated buffer of global memory.
		this->setBufferSize(numthreads, size);
	}

	// Set the size of the pre-allocated buffer of global memory.
	template<typename T>
	void DynamicGlobalMemory<T>::setBufferSize(int numthreads, int size)
	{
#ifdef _DEBUG
		std::cout << "void DynamicGlobalMemory<T>::setBufferSize(int numthreads, int size)" << std::endl;
#endif
		// Set the dimension size.
		DimensionSizeVector dim(2);
		dim[0] = numthreads;
		dim[1] = size;

		// Set the number of threads.
		this->numthreads = numthreads;

		// Set the number of dynamic allocations available to each thread.
		this->size = size;

		// Allocate buffer.
		if(!this->buffer.allocateMemory(dim)) {
			std::cout << "Failed to allocate dynamic memory" << std::endl;
		}
	}

	// // Set the indexes to 0 for each thread.
	template<typename T>
	void DynamicGlobalMemory<T>::resetIndex()
	{
#ifdef _DEBUG
		std::cout << "void DynamicGlobalMemory<T>::resetIndex()" << std::endl;
#endif
		// Set the indexes to 0.
		cudaMemset(this->currentIndex, 0, sizeof(int) * this->numthreads);
	}

	// Callable only on the device.
	// Returns a specified size of free memory in the buffer to a thread.
	template<typename T>
	__device__ T* DynamicGlobalMemory<T>::create(int thread, int length = 1)
	{
		// Specify index.
		int i[] = { this->currentIndex[thread], thread};

		// Set index to the next free available position in memory.
		this->currentIndex[thread] += length;

		// Return a pointer to the free memory.
		return &this->buffer.memory[this->buffer.getIndex(i)];
	}
}
